#include <stdio.h>
#include <stdlib.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <sys/fcntl.h>
#include <sys/mman.h>
#include <unistd.h>
#include <string.h>
#include <CL/cl.h>
#include <time.h>
#include <string>
#include <malloc.h>
#include "common.h"
#include "scanImpl.cpp"
#include "../../kernel/para.h"

using namespace std;
struct statistic pp;
struct clContext context;


#define BI 1000000000
#define HSIZE (1024*1024*16)

extern const char * createProgram(string, int *);

int main(int argc, char * argv[]){
	int *fact, *dim;
	int dimSize, factSize;
	int factNum, dimNum;

	int *cpu_count, *resPsum;
	int count = 0;

	struct timespec start,end, afBuild;
	char **factTable, **dimTable;

	if(argc!=5){
		printf("error: factTable dimTable fColNum dColNum\n");
		exit(-1);
	}

	int fColNum = atoi(argv[3]);
	int dColNum = atoi(argv[4]);

	fact =(int *) malloc(sizeof(int) * fColNum);
	dim = (int *) malloc(sizeof(int) * dColNum);
	dimTable = (char **) malloc(sizeof(char *) * dColNum);
	factTable = (char **) malloc(sizeof(char *) * fColNum);

	if(!fact || !dim ){
		printf("Failed to allocate host memory\n");
		exit(-1);
	}

	for(int i=0;i<fColNum;i++){
		char buf[32] = {0};
		sprintf(buf,"%s%d",argv[1],i);
		fact[i] = open(buf,O_RDONLY);
	}
	for(int i = 0;i<dColNum;i++){
		char buf[32] = {0};
		sprintf(buf,"%s%d",argv[2],i);
		dim[i] = open(buf, O_RDONLY);
	}

	dimSize = lseek(dim[0],0,SEEK_END);
	factSize = lseek(fact[0],0,SEEK_END);
	factNum = factSize/sizeof(int);
	dimNum = dimSize/sizeof(int);

	for(int i=0;i<fColNum;i++){
		factTable[i] = (char *) memalign(128,factSize);
		char *factT = (char *)mmap(0,factSize,PROT_READ,MAP_SHARED,fact[i],0);
		memcpy(factTable[i],factT,factSize);
		munmap(factT,factSize);
	}
	for(int i=0;i<dColNum;i++){
		dimTable[i] = (char *) memalign(128,dimSize);
		char *dimT = (char *)mmap(0,dimSize,PROT_READ,MAP_SHARED,dim[i],0);
		memcpy(dimTable[i],dimT,dimSize);
		munmap(dimT,dimSize);
	}

	int psc = 0;
	const char *ps = createProgram("../kernel/join.cl",&psc);

/*
 * 	build hash table on GPU
 */
	
	clock_gettime(CLOCK_REALTIME,&start);

	size_t localSize = LOCAL;
	size_t globalSize = GLOBAL * localSize;
	size_t threadNum = globalSize;
	cl_int error = 0;
        cl_device_id device;

        cl_uint numP;
        clGetPlatformIDs(0,NULL,&numP);

        cl_platform_id * pid = new cl_platform_id[numP];
        clGetPlatformIDs(numP, pid, NULL);

        clGetDeviceIDs(pid[PID],DEVICETYPE,1,&device,NULL);

        context.context = clCreateContext(0,1,&device,NULL,NULL,&error);
        if(error != CL_SUCCESS){
                printf("Error creating context\n");
                exit(-1);
        }

        context.queue = clCreateCommandQueue(context.context, device, 0, &error);
        if(error != CL_SUCCESS){
                printf("Error creating command queue\n");
                exit(-1);
        }

        context.program = clCreateProgramWithSource(context.context, psc, (const char **)&ps, 0, &error);
        if(error != CL_SUCCESS){
                printf("Error creating program\n");
                exit(-1);
        }

	error = clBuildProgram(context.program, 0, 0 , 0 , 0, 0);
        if(error != CL_SUCCESS){
                printf("Error building program\n");
                exit(-1);
        }

	clock_gettime(CLOCK_REALTIME,&afBuild);

	int tnum = HSIZE;
	cl_mem gpu_hashNum = clCreateBuffer(context.context, CL_MEM_READ_WRITE, sizeof(int)*HSIZE, NULL, &error);
	context.kernel = clCreateKernel(context.program, "cl_memset_int",0);
	clSetKernelArg(context.kernel,0,sizeof(cl_mem), (void *)&gpu_hashNum);
	clSetKernelArg(context.kernel,1,sizeof(int),(void *)&tnum );
	clEnqueueNDRangeKernel(context.queue,context.kernel, 1,0,&globalSize,&localSize,0,0,0);

	cl_mem gpuFilter = clCreateBuffer(context.context, CL_MEM_READ_WRITE, factSize, NULL, &error);

	cl_mem gpu_count = clCreateBuffer(context.context, CL_MEM_READ_WRITE, sizeof(int)*threadNum, NULL, &error);
	cl_mem gpu_result = clCreateBuffer(context.context, CL_MEM_WRITE_ONLY, factNum * sizeof(int), NULL, &error);

	cl_mem gpu_psum = clCreateBuffer(context.context, CL_MEM_READ_WRITE, sizeof(int)*HSIZE,NULL,&error);

	cl_mem gpu_psum1 = clCreateBuffer(context.context, CL_MEM_READ_WRITE, sizeof(int)*HSIZE,NULL,&error);

	cl_mem gpu_bucket = clCreateBuffer(context.context, CL_MEM_READ_WRITE, 2*dimSize, NULL, &error);

	cl_mem gpu_resPsum = clCreateBuffer(context.context, CL_MEM_READ_WRITE, sizeof(int) * threadNum, NULL, &error);

	cl_mem gpu_dim = clCreateBuffer(context.context, CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR,dimSize,dimTable[0],&error);
	//clEnqueueWriteBuffer(context.queue,gpu_dim,CL_TRUE,0,dimSize,dimTable[0],0,0,0);

	context.kernel = clCreateKernel(context.program,"count_hash_num",0);

	clSetKernelArg(context.kernel,0,sizeof(cl_mem),(void *)&gpu_dim);
	clSetKernelArg(context.kernel,1,sizeof(int), (void *)&dimNum);
	clSetKernelArg(context.kernel,2,sizeof(cl_mem),(void *)&gpu_hashNum);

	clEnqueueNDRangeKernel(context.queue,context.kernel, 1,0,&globalSize,&localSize,0,0,0);

	scanImpl(gpu_hashNum,HSIZE,gpu_psum,&context, &pp);

	clEnqueueCopyBuffer(context.queue,gpu_psum,gpu_psum1,0,0,sizeof(int)*HSIZE,0,0,0);

	context.kernel = clCreateKernel(context.program,"build_hash_table",0);

	clSetKernelArg(context.kernel,0,sizeof(cl_mem), (void *)&gpu_dim);
	clSetKernelArg(context.kernel,1,sizeof(int), (void *) &dimNum);
	clSetKernelArg(context.kernel,2,sizeof(cl_mem), (void *)&gpu_psum1);
	clSetKernelArg(context.kernel,3,sizeof(cl_mem), (void *)&gpu_bucket);
	clEnqueueNDRangeKernel(context.queue,context.kernel, 1,0,&globalSize,&localSize,0,0,0);


/*
 *	join on GPU
*/


	cl_mem gpu_fact = clCreateBuffer(context.context, CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR, factSize, factTable[0], &error);
	//clEnqueueWriteBuffer(context.queue,gpu_fact,CL_TRUE,0,factSize,factTable[0],0,0,0);

	count = HSIZE;

	context.kernel = clCreateKernel(context.program,"count_join_result",0);

	clSetKernelArg(context.kernel,0,sizeof(cl_mem), (void *)&gpu_hashNum);
	clSetKernelArg(context.kernel,1,sizeof(cl_mem), (void *)&gpu_psum);
	clSetKernelArg(context.kernel,2,sizeof(cl_mem), (void *)&gpu_bucket);
	clSetKernelArg(context.kernel,3,sizeof(cl_mem), (void *)&gpu_fact);
	clSetKernelArg(context.kernel,4,sizeof(int), (void *)&factNum);
	clSetKernelArg(context.kernel,5,sizeof(cl_mem), (void *)&gpu_count);
	clSetKernelArg(context.kernel,6,sizeof(cl_mem), (void *)&gpuFilter);

	clEnqueueNDRangeKernel(context.queue,context.kernel, 1,0,&globalSize,&localSize,0,0,0);

	cpu_count = new int[threadNum];
	clEnqueueReadBuffer(context.queue, gpu_count, CL_TRUE, 0, sizeof(int) * threadNum, cpu_count,0,0,0);

	resPsum = new int[threadNum]; 

	scanImpl(gpu_count,threadNum,gpu_resPsum, &context,&pp);

	clEnqueueReadBuffer(context.queue, gpu_resPsum, CL_TRUE, 0, sizeof(int) * threadNum, resPsum,0,0,0);

	count = resPsum[threadNum-1] + cpu_count[threadNum-1];

	printf("%d\n",count);

	for(int i=1;i<fColNum;i++){
		gpu_fact = clCreateBuffer(context.context,CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR, factSize, factTable[i],&error);
		//clEnqueueWriteBuffer(context.queue,gpu_fact,CL_TRUE,0,factSize,factTable[i],0,0,0);

		context.kernel = clCreateKernel(context.program,"joinFact",0);
		int tmp = sizeof(int);
		clSetKernelArg(context.kernel,0,sizeof(cl_mem), (void *)&gpu_resPsum);
		clSetKernelArg(context.kernel,1,sizeof(cl_mem), (void *)&gpu_fact);
		clSetKernelArg(context.kernel,2,sizeof(int), (void *)&tmp);
		clSetKernelArg(context.kernel,3,sizeof(int), (void *)&factNum);
		clSetKernelArg(context.kernel,4,sizeof(cl_mem), (void *)&gpuFilter);
		clSetKernelArg(context.kernel,5,sizeof(cl_mem), (void *)&gpu_result);
		
		clEnqueueNDRangeKernel(context.queue,context.kernel, 1,0,&globalSize,&localSize,0,0,0);
		clEnqueueMapBuffer(context.queue,gpu_result,CL_TRUE,CL_MAP_READ,0,sizeof(int)*count, 0,0,0,0);
	}

	for(int i=1; i<dColNum;i++){
		gpu_dim = clCreateBuffer(context.context,CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR, dimSize, dimTable[i],&error);
		//clEnqueueWriteBuffer(context.queue,gpu_dim,CL_TRUE,0,dimSize,dimTable[i],0,0,0);
		context.kernel = clCreateKernel(context.program,"joinDim",0);
		int tmp = sizeof(int);
		clSetKernelArg(context.kernel,0,sizeof(cl_mem), (void *)&gpu_resPsum);
		clSetKernelArg(context.kernel,1,sizeof(cl_mem), (void *)&gpu_dim);
		clSetKernelArg(context.kernel,2,sizeof(int), (void *)&tmp);
		clSetKernelArg(context.kernel,3,sizeof(int), (void *)&factNum);
		clSetKernelArg(context.kernel,4,sizeof(cl_mem), (void *)&gpuFilter);
		clSetKernelArg(context.kernel,5,sizeof(cl_mem), (void *)&gpu_result);
		
		clEnqueueNDRangeKernel(context.queue,context.kernel, 1,0,&globalSize,&localSize,0,0,0);
		clEnqueueMapBuffer(context.queue,gpu_result,CL_TRUE,CL_MAP_READ,0,sizeof(int)*count, 0,0,0,0);
	}
	clFinish(context.queue);
	clReleaseCommandQueue(context.queue);

	clock_gettime(CLOCK_REALTIME,&end);
	double timeE = (end.tv_sec -  start.tv_sec)* BI + end.tv_nsec - start.tv_nsec;
	printf("CPU time total: %lf\n", timeE/1000/1000);
	timeE = (end.tv_sec -  afBuild.tv_sec)* BI + end.tv_nsec - afBuild.tv_nsec;
	printf("CPU time without build: %lf\n", timeE/1000/1000);

	for(int i=0;i<fColNum;i++){
		close(fact[i]);
		free(factTable[i]);
	}
	for(int i = 0;i<dColNum;i++){
		close(dim[i]);
		free(dimTable[i]);
	}

	return 0;
}
